#include "cutil.h"
#include "zspace.h"
#include "wsecondspaceapp.h"

/**
    @file zspace.cpp
    @author Claus Wimmer
    @date 2014-01-04

    Project: secondspace
    Licence: LGPL (Lesser GPL)
*/

const char* ZSpace::_kern_init_prog = "#include \"kern_include.cl\"\n#include \"kern_init.cl\"\n";
const char* ZSpace::_kern_step_prog = "#include \"kern_include.cl\"\n#include \"kern_step.cl\"\n";

ZSpace::ZSpace(QWidget *parent) :
    // Begin of member init.
    QObject(parent),
    _imgData(new cl_uchar[(size_t)(4 * points() / CUtil::layers)]),
    _image(QImage((const cl_uchar*)(this->_imgData), CUtil::pixwidth, CUtil::pixheight, QImage::Format_RGB32)),
    _spaceData(new cl_float[points() * CUtil::varsiz]),
    _spaceBits(new cl_int[points()]),
    _errNum(0),
    _slider(CUtil::valslider),
    _idx(0),
    _rows(CUtil::pixheight),
    _cols(CUtil::pixwidth),
    _layers(CUtil::layers),
    _varsiz(CUtil::varsiz),
    _step(0),
    _mouse(new cl_uint[2]),
    _global_work_group_size(points()),
    _local_work_group_size(0),
    _deviceIDs(0),
    _context(0)
    // End of member init.
{
    // Mouse.
    _mouse[MOUSEX] = CUtil::pixwidth / 2;
    _mouse[MOUSEY] = CUtil::pixheight / 2;
    // Platform.
    _errNum = clGetPlatformIDs(0, 0, &_numPlatforms);
    checkErr(_errNum, "clGetPlatformIDs / 1");
    CUtil::showMessage(QString("Num OpenCL Platforms = %1.").arg(_numPlatforms));
    _platformIds = new cl_platform_id[_numPlatforms];
    assert(0 != _platformIds);
    _errNum = clGetPlatformIDs(_numPlatforms, _platformIds, 0);
     checkErr(_errNum, "clGetPlatformIds / 2");
    // Device.
    for(; _idx < _numPlatforms; ++_idx){
        _errNum = clGetDeviceIDs(_platformIds[_idx], CL_DEVICE_TYPE_GPU, 0, 0, &_numDevices);
        checkErr(_errNum, "clGetDeviceIDs / 1");
        _deviceIDs = new cl_device_id[_numDevices];
        _errNum = clGetDeviceIDs(_platformIds[_idx], CL_DEVICE_TYPE_GPU, _numDevices, _deviceIDs, 0);
        checkErr(_errNum, "clGetDeviceIDs / 2");
        CUtil::showMessage(QString("Number of OpenCL devices = %1.").arg(_numDevices));
        if(0 < _numDevices){
            break;
        }
    }
    if(0 == _numDevices){
        checkErr(-1, "No GPU device found.");
    }
    for(cl_uint i = 0; i < _numDevices; ++i){
        size_t device_global_mem_size = -1, device_max_mem_alloc_size = -1;
        _errNum = clGetDeviceInfo(_deviceIDs[i], CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof(_max_work_group_size), &_max_work_group_size, NULL);
        CUtil::showMessage(QString("Device[%1] of platform[%2] has %3 parallel kernel instances (CL_DEVICE_MAX_WORK_GROUP_SIZE).").arg(i).arg(_idx).arg(_max_work_group_size));
        _errNum |= clGetDeviceInfo(_deviceIDs[i], CL_DEVICE_GLOBAL_MEM_SIZE, sizeof(size_t), &device_global_mem_size, NULL);
        CUtil::showMessage(QString("Device[%1] of platform[%2] can use memory with a size of %3 Mbytes (CL_DEVICE_GLOBAL_MEM_SIZE).").arg(i).arg(_idx).arg(device_global_mem_size / 1024 / 1024));
        _errNum |= clGetDeviceInfo(_deviceIDs[i], CL_DEVICE_MAX_MEM_ALLOC_SIZE, sizeof(size_t), &device_max_mem_alloc_size, NULL);
        CUtil::showMessage(QString("Device[%1] of platform[%2] can alloc memory with a size of %3 Mbytes (CL_DEVICE_MAX_MEM_ALLOC_SIZE).").arg(i).arg(_idx).arg(device_max_mem_alloc_size / 1024 / 1024));
    }
    // Context.
    cl_context_properties _contextProperties[] = {CL_CONTEXT_PLATFORM, (cl_context_properties)_platformIds[_idx], 0};
    _context = clCreateContext(_contextProperties, _numDevices, _deviceIDs, &ZSpace::contextCallback, 0, &_errNum);
    checkErr(_errNum, "clCreateContext / 1");
    // Kern - read / init.
    _prog_init = clCreateProgramWithSource(_context, 1, &ZSpace::_kern_init_prog, 0, &_errNum);
    checkErr(_errNum, "clCreateProgramWithSource / kern_init");
    _errNum = clBuildProgram(_prog_init, _numDevices, _deviceIDs, CUtil::includepath.toLatin1().constData(), 0, 0);
    if(CL_SUCCESS != _errNum) {
        clGetProgramBuildInfo(_prog_init, _deviceIDs[0], CL_PROGRAM_BUILD_LOG, sizeof(_buildlog), _buildlog, 0);
        CUtil::showMessage(QString("Error in init-kernel : \n%1\n").arg(_buildlog));
        checkErr(_errNum, "clBuildProgram / kern_init");
    }
    _kern_init = clCreateKernel(_prog_init, "kern_init", &_errNum);
    checkErr(_errNum, "clCreateKernel / kern_init");
    // Kern - read / step.
    _prog_step = clCreateProgramWithSource(_context, 1, &ZSpace::_kern_step_prog, 0, &_errNum);
    checkErr(_errNum, "clCreateProgramWithSource / kern_step");
    _errNum = clBuildProgram(_prog_step, _numDevices, _deviceIDs, CUtil::includepath.toLatin1().constData(), 0, 0);
    if(CL_SUCCESS != _errNum) {
        clGetProgramBuildInfo(_prog_step, _deviceIDs[0], CL_PROGRAM_BUILD_LOG, sizeof(_buildlog), _buildlog, 0);
        CUtil::showMessage(QString("Error in step-kernel : \n%1\n").arg(_buildlog));
        checkErr(_errNum, "clBuildProgram / kern_step");
    }
    _kern_step = clCreateKernel(_prog_step, "kern_step", &_errNum);
    checkErr(_errNum, "clCreateKernel / kern_step");
    // Copy memory buffer.
    _mem_spaceData =    clCreateBuffer(_context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, (size_t)(sizeof(cl_float) * points() * CUtil::varsiz), (void*)_spaceData, &_errNum);
    checkErr(_errNum, "clCreateBuffer / _mem_spaceData");
    _mem_spaceBits =    clCreateBuffer(_context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, (size_t)(sizeof(cl_uint ) * points()        ), (void*)_spaceBits     , &_errNum);
    checkErr(_errNum, "clCreateBuffer / _mem_spaceBits");
    _mem_imgData =      clCreateBuffer(_context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, (size_t)(sizeof(cl_uchar ) * points() / CUtil::layers * 4), (void*)_imgData, &_errNum);
    checkErr(_errNum, "clCreateBuffer / _imgData");
    _mem_step =         clCreateBuffer(_context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, (size_t)(sizeof(cl_uint  )                   ), (void*)&_step        , &_errNum);
    checkErr(_errNum, "clCreateBuffer / _step");
    _mem_slider =       clCreateBuffer(_context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, (size_t)(sizeof(cl_int   )                   ), (void*)&_slider      , &_errNum);
    checkErr(_errNum, "clCreateBuffer / _slider");
    _mem_mouse  =       clCreateBuffer(_context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, (size_t)(sizeof(cl_uint  ) * 2               ), (void*)_mouse        , &_errNum);
    checkErr(_errNum, "clCreateBuffer / _mouse");
    _mem_modus  =       clCreateBuffer(_context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, (size_t)(sizeof(cl_uint  ) * 1               ), (void*)&_modus       , &_errNum);
    checkErr(_errNum, "clCreateBuffer / _modus");
    // Card queue.
    _queue = clCreateCommandQueue(_context, _deviceIDs[0], 0, &_errNum);
    checkErr(_errNum, "clCreateCommandQueue");
    // Init kern args.
    _errNum = 0;
    _errNum |= clSetKernelArg(_kern_init, 0 , sizeof(cl_mem), &_mem_spaceData);
    _errNum |= clSetKernelArg(_kern_init, 1 , sizeof(cl_mem), &_mem_spaceBits);
    _errNum |= clSetKernelArg(_kern_init, 2 , sizeof(cl_mem), &_mem_imgData);
    _errNum |= clSetKernelArg(_kern_init, 3 , sizeof(cl_mem), &_mem_step);
    _errNum |= clSetKernelArg(_kern_init, 4 , sizeof(cl_mem), &_mem_slider);
    _errNum |= clSetKernelArg(_kern_init, 5 , sizeof(cl_mem), &_mem_mouse);
    _errNum |= clSetKernelArg(_kern_init, 6 , sizeof(cl_mem), &_mem_modus);
    _errNum |= clSetKernelArg(_kern_init, 7 , sizeof(cl_uint), &_rows);
    _errNum |= clSetKernelArg(_kern_init, 8 , sizeof(cl_uint), &_cols);
    _errNum |= clSetKernelArg(_kern_init, 9 , sizeof(cl_uint), &_layers);
    _errNum |= clSetKernelArg(_kern_init, 10, sizeof(cl_uint), &_varsiz);
    checkErr(_errNum, "clSetKernelArg / kern_init");
    //Step kern args.
    _errNum = 0;
    _errNum |= clSetKernelArg(_kern_step, 0 , sizeof(cl_mem), &_mem_spaceData);
    _errNum |= clSetKernelArg(_kern_step, 1 , sizeof(cl_mem), &_mem_spaceBits);
    _errNum |= clSetKernelArg(_kern_step, 2 , sizeof(cl_mem), &_mem_imgData);
    _errNum |= clSetKernelArg(_kern_step, 3 , sizeof(cl_mem), &_mem_step);
    _errNum |= clSetKernelArg(_kern_step, 4 , sizeof(cl_mem), &_mem_slider);
    _errNum |= clSetKernelArg(_kern_step, 5 , sizeof(cl_mem), &_mem_mouse);
    _errNum |= clSetKernelArg(_kern_step, 6 , sizeof(cl_mem), &_mem_modus);
    _errNum |= clSetKernelArg(_kern_step, 7 , sizeof(cl_uint), &_rows);
    _errNum |= clSetKernelArg(_kern_step, 8 , sizeof(cl_uint), &_cols);
    _errNum |= clSetKernelArg(_kern_step, 9 , sizeof(cl_uint), &_layers);
    _errNum |= clSetKernelArg(_kern_step, 10, sizeof(cl_uint), &_varsiz);
    checkErr(_errNum, "clSetKernelArg / kern_step");
}

ZSpace::~ZSpace() {
    cleanup();
}

void ZSpace::initSpace() {

    _step = 0;
    _modus = CUtil::modus;
    // Copy the input data of the kernel.
    _errNum = clEnqueueWriteBuffer(_queue, _mem_spaceBits, CL_TRUE, 0, sizeof(cl_uint) * points(), (void*)_spaceBits, 0, 0, 0);
    checkErr(_errNum, "clEnqueueWriteBuffer / kern_init / spaceBits");
    _errNum = clEnqueueWriteBuffer(_queue, _mem_step, CL_TRUE, 0, sizeof(cl_uint), (void*)&_step, 0, 0, 0);
    checkErr(_errNum, "clEnqueueWriteBuffer / kern_init / step");
    _errNum = clEnqueueWriteBuffer(_queue, _mem_slider, CL_TRUE, 0, sizeof(cl_int), (void*)&_slider, 0, 0, 0);
    checkErr(_errNum, "clEnqueueWriteBuffer / kern_init / slider");
    _errNum = clEnqueueWriteBuffer(_queue, _mem_mouse, CL_TRUE, 0, (size_t)(sizeof(cl_int) * 2), (void*)_mouse, 0, 0, 0);
    checkErr(_errNum, "clEnqueueWriteBuffer / kern_init / mouse");
    _errNum = clEnqueueWriteBuffer(_queue, _mem_modus, CL_TRUE, 0, (size_t)(sizeof(cl_int) * 1), (void*)&_modus, 0, 0, 0);
    checkErr(_errNum, "clEnqueueWriteBuffer / kern_init / modus");
    _errNum = clEnqueueNDRangeKernel(_queue, _kern_init, 1, 0, &_global_work_group_size, 0, 0, 0, 0);
    checkErr(_errNum, "clEnqueueNDRangeKernel / kern_init");
    // Exec the init kernel.
    _errNum = clEnqueueReadBuffer(_queue, _mem_imgData, CL_TRUE, 0, (size_t)(sizeof(cl_uchar) * points() / CUtil::layers * 4), _imgData, 0, 0, 0);
    checkErr(_errNum, "clEnqueueReadBuffer / kern_init / imgData");
    emit imageReady();

}

void ZSpace::stepSpace() {
    // Exec the step kernel four times. This is equal to 2 graphic images.
    // Depending on the preferences only each n-th picture is shown on the screen.
    for(int _idx = 0; _idx < CUtil::numsteps; ++_idx){
        _modus = CUtil::modus;
        // Increase the counter.
        ++_step;
        // Copy the input data of the kernel.
        _errNum = clEnqueueWriteBuffer(_queue, _mem_step, CL_TRUE, 0, sizeof(cl_uint), (void*)&_step, 0, 0, 0);
        checkErr(_errNum, "clEnqueueWriteBuffer / kern_step / step");
        _errNum = clEnqueueWriteBuffer(_queue, _mem_slider, CL_TRUE, 0, sizeof(cl_int), (void*)&_slider, 0, 0, 0);
        checkErr(_errNum, "clEnqueueWriteBuffer / kern_step / slider");
        // Execute the step kernel.
        _errNum = clEnqueueWriteBuffer(_queue, _mem_mouse, CL_TRUE, 0, (size_t)(sizeof(cl_int) * 2), (void*)_mouse, 0, 0, 0);
        checkErr(_errNum, "clEnqueueWriteBuffer / kern_step / mouse");
        _errNum = clEnqueueWriteBuffer(_queue, _mem_modus, CL_TRUE, 0, (size_t)(sizeof(cl_int) * 1), (void*)&_modus, 0, 0, 0);
        checkErr(_errNum, "clEnqueueWriteBuffer / kern_step / modus");
        // Copy the graphical image from the kernel to show it.
        _errNum = clEnqueueNDRangeKernel(_queue, _kern_step, 1, 0, &_global_work_group_size, 0, 0, 0, 0);
    }
    _errNum = clEnqueueReadBuffer(_queue, _mem_imgData, CL_TRUE, 0, (size_t)(sizeof(cl_uchar) * points() / CUtil::layers * 4), _imgData, 0, 0, 0);
    checkErr(_errNum, "clEnqueueReadBuffer / kern_step / imgData");
    emit imageReady();
}

// Sent from the main application window.
void ZSpace::sliderValueChanged(int value) {
    this->_slider = value;
}

// Width times height from pixel space.
uint ZSpace::points() const {
    return CUtil::pixwidth * CUtil::pixheight * CUtil::layers;
}

// In case of abnormal termination.
void ZSpace::checkErr(cl_int _err, const char *_name){
    if(CL_SUCCESS != _err){
        QString msg = QString("Graphic Card mass calculator OPENCL raised error = %1 (number = %2). EXIT!").arg(_name).arg(_err);
        // Console output.
        CUtil::showMessage(msg);
        // Popup window output.
        QMessageBox box(
            QMessageBox::Critical,
            "Secondspace - crash",
            msg,
            QMessageBox::Ok,
            (QWidget*)parent()
        );
        // Stylsheet specific to the popup window.
        box.setStyleSheet(CUtil::qss());
        box.exec();
        // Return abnormal termination number.
        exit(-1);
    }
}

// Only normal termination covered.
void ZSpace::cleanup() {

    if(0 !=_mem_spaceData) {
        clReleaseMemObject(_mem_spaceData);
    }
    if(0 !=_mem_spaceBits) {
        clReleaseMemObject(_mem_spaceBits);
    }
    if(0 !=_mem_imgData) {
        clReleaseMemObject(_mem_imgData);
    }
    if(0 !=_mem_step) {
        clReleaseMemObject(_mem_step);
    }
    if(0 !=_mem_slider) {
        clReleaseMemObject(_mem_slider);
    }
    if(0 !=_mem_mouse) {
        clReleaseMemObject(_mem_mouse);
    }
    if(0 != _queue) {
        clReleaseCommandQueue(_queue);
    }
    if(0 != _kern_init) {
        clReleaseKernel(_kern_init);
    }
    if(0 != _kern_step) {
        clReleaseKernel(_kern_step);
    }
    if(0 != _prog_init) {
        clReleaseProgram(_prog_init);
    }
    if(0 != _prog_step) {
        clReleaseProgram(_prog_step);
    }
    if(0 != _context) {
        clReleaseContext(_context);
    }

    delete[] _imgData;
    delete _deviceIDs;
    delete _platformIds;
    delete[] _spaceBits;
    delete[] _spaceData;
    delete[] _mouse;

}

// GPU error messages.
void CL_CALLBACK ZSpace::contextCallback(const char *_errInfo, const void *_privateInfo, size_t, void*){
    CUtil::showMessage(QString("Error occured while processing mass calculation with OpenCL, info =  %1, %2.").arg(_errInfo).arg((char*)_privateInfo));
}

QImage *ZSpace::image() {
    return &_image;
}

// Sent from main application window.
void ZSpace::mousePointed(uint x, uint y) {
    _mouse[MOUSEX] = x;
    _mouse[MOUSEY] = y;
}
